当SYCL实现构建在另一个API上时，该实现可能能够与使用底层API机制定义的内核进行互操作。这允许应用程序将SYCL集成到现有代码库中。\par

因为SYCL实现可能基于许多API，所以本节中描述的功能是可选的，并不是所有实现都支持。底层API甚至可能根据特定的设备类型，或设备供应商而有所不同!\par

广义地说，一个实现可能支持两种互操作性机制:从API定义的源或中间表示(IR)或从特定于API的句柄。这两种机制中，从API定义的源或中间表示创建内核的能力更具可移植性，因为一些源或IR格式由多个API支持。例如，OpenCL C内核可直接使用，或者可以编译成某个API可以理解的形式，但是来自某个API的特定于API的内核句柄，不太可能让另一个API理解。\par

\begin{tcolorbox}[colback=red!5!white,colframe=red!75!black]
记住，所有形式的操作都是可选的!不同的SYCL实现可能支持从不同API句柄创建的内核——或者根本不支持。要了解详细信息，请查阅文档!
\end{tcolorbox}

\hspace*{\fill} \par %插入空行
\textbf{与API定义的互动性}

使用这种形式的互动性，内核的内容描述为源代码或使用不是由SYCL定义的中间表示，但是内核对象仍然使用SYCL API创建。这种形式的互动性可以重用，使用其他源语言编写的内核库，或者使用中间表示形式生成代码的特定语言(DSl)。\par

实现必须理解内核源代码或中间表示，才能利用这种形式的互动性。例如，如果内核以源码的形式使用OpenCL C编写，那么实现必须支持从OpenCL C内核源代码构建SYCL程序。\par

图10-7展示了如何将SYCL内核编写为OpenCL C内核源代码。\par

\hspace*{\fill} \par %插入空行
图10-7 使用OpenCL C内核源码创建SYCL内核
\begin{lstlisting}[caption={}]
// Note: This must select a device that supports interop!
queue Q{ cpu_selector{} };

program p{Q.get_context()};

p.build_with_source(R"CLC(
		kernel void add(global int* data) {
			int index = get_global_id(0);
			data[index] = data[index] + 1;
		}
		)CLC",
		"-cl-fast-relaxed-math");
		
std::cout << "Running on device: "
		  << Q.get_device().get_info<info::device::name>() << "\n";

Q.submit([&](handler& h) {
	accessor data_acc {data_buf, h};
	
	h.set_args(data_acc);
	h.parallel_for(size, p.get_kernel("add"));
});
\end{lstlisting}

例子中，内核以字符串形式表示，在同一个文件中对SYCL的主机API进行调用(但这并不是必须的)，一些程序可以从文件读取内核字符串然后生成内核。\par

因为SYCL编译器无法看到用内核，所以必须使用set\_arg()或set\_args()接口显式地传递内核参数。SYCL运行时和API定义语言必须将对象作为内核参数。本例中，访问器dataAcc作为内核的全局指针参数进行数据传递。\par

build\_with\_source()接口支持传递API定义的构建选项，来控制内核的编译方式。例子中，编译选项为-cl-fast- relaxation -math，用于表示内核编译器可以使用更快、低精度的数学库。编译选项是可选的，如果不需要生成选项，可以省略。。\par

\hspace*{\fill} \par %插入空行
\textbf{与API定义内核的互动性}

内核对象本身在另一个API中创建，然后导入SYCL。这种形式的互动性使，应用程序的一部分可以使用底层API直接创建和使用内核对象，而应用程序的另一部分可以使用SYCL API重用相同的内核。图10-8中的代码展示了如何从OpenCL内核对象创建SYCL内核。\par

\hspace*{\fill} \par %插入空行
图10-8 OpenCL内核对象创建的内核
\begin{lstlisting}[caption={}]
// Note: This must select a device that supports interop 
// with OpenCL kernel objects!
queue Q{ cpu_selector{} };
context sc = Q.get_context();

const char* kernelSource =
	R"CLC(
	kernel void add(global int* data) {
		int index = get_global_id(0);
		data[index] = data[index] + 1;
	}
	)CLC";
cl_context c = sc.get();
cl_program p =
	clCreateProgramWithSource(c, 1, &kernelSource, nullptr, nullptr);
clBuildProgram(p, 0, nullptr, nullptr, nullptr, nullptr);
cl_kernel k = clCreateKernel(p, "add", nullptr);

std::cout << "Running on device: "
	 	  << Q.get_device().get_info<info::device::name>() << "\n";
	 	  
Q.submit([&](handler& h) {
	accessor data_acc{data_buf, h};
	
	h.set_args(data_acc);
	h.parallel_for(size, kernel{k, sc});
});

clReleaseContext(c);
clReleaseProgram(p);
clReleaseKernel(k);
\end{lstlisting}

与其他形式的互动性一样，SYCL编译器对API定义的内核对象不可见。因此，必须使用set\_arg()或set\_args()接口显式传递参数，并且SYCL运行时和底层API必须遵循传递参数的约定。\par









































